__global__ void kernel_euc_norm(float *inp, float *out, size_t elems_per_thread)
{
    // 快速定义DTE
    tops_dte_ctx_t ctx;
    tops::dte_scope s(ctx);

    auto thread_idx = 0;
    static const size_t L1_BUFFER_LEN = 1280;
    auto pos_start = 0;
    auto pos_end = elems_per_thread;

    __valigned__ float l1_buffer_input[L1_BUFFER_LEN];
    __valigned__ float l1_buffer_output[L1_BUFFER_LEN];

    tops::mdspan l1_inp(tops::Private, &l1_buffer_input, L1_BUFFER_LEN);
    tops::mdspan l1_out(tops::Private, &l1_buffer_output, L1_BUFFER_LEN);
    tops::mdspan l3_out(tops::Global, out, 1);

    float sum_num = 0; // 用于存储平方和

    for (size_t i = pos_start; i < pos_end; i += L1_BUFFER_LEN)
    {
        tops::mdspan l3_inp(tops::Global, inp + i, L1_BUFFER_LEN);

        // L3 to L1
        tops::memcpy(ctx, l1_inp, l3_inp);

        for (size_t j = 0; j < L1_BUFFER_LEN; j += tops::vlength<vfloat>())
        {
            const auto &v_inp = tops::vload<vfloat>(l1_buffer_input + j);
            // 计算平方和
            const auto &v_pow = tops::vmul<vfloat>(v_inp, v_inp);
            for (int k = 0; k < tops::vlength<vfloat>(); k++)
            {
                sum_num += v_pow[k];
            }
        }
    }

    // 计算欧几里得范数
    sum_num = sqrtf(sum_num);
    tops::mdspan p(tops::Private, &sum_num, 1);
    // L1 to L3
    tops::memcpy(ctx, l3_out, sum_num);
}


void GCU_EUC_NORM(float *__restrict dev_inp, float *__restrict dev_out, const int nr_elems)
{
    static const size_t blocks = 1;
    static const size_t threads = 1;
    kernel_euc_norm<<<blocks, threads>>>(dev_inp, dev_out, nr_elems);
}